// Csharp_C.cpp : Defines the exported functions for the DLL application.
//
#include "stdafx.h"
#include "Csharp_C.h"
#include <boost/graph/adjacency_list.hpp>
#include <boost/graph/small_world_generator.hpp>
#include <boost/random/linear_congruential.hpp>
#include <boost/graph/graph_traits.hpp>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

#define MEM_SIZE (128)
#define MAX_SOURCE_SIZE (0x100000)


#define MAX_ITEM_PER_BLOCK 512


//Structure to hold a node information
struct Node
{
	int starting;
	int no_of_edges;
};

typedef boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS> Graph;
typedef boost::small_world_iterator<boost::minstd_rand, Graph> SWGen;
size_t no_of_nodes = 0;
boost::graph_traits<Graph>::edges_size_type edge_list_size = 0;


// This is an example of an exported variable
CSHARP_C_API int nCsharp_C = 0;

// This is an example of an exported function.
extern "C" {
	CSHARP_C_API int fnCsharp_C(void)
	{
		/// Graph Generation
		boost::minstd_rand gen;
		// Create graph with 100 nodes 
		no_of_nodes = 10000;

		std::cout << "Generation of Small World Graph is started....." << std::endl;
		Graph* g = new Graph(SWGen(gen, no_of_nodes, 6, 0.03), SWGen(), no_of_nodes);
		std::cout << "Generation of Small World Graph has been finished....." << std::endl;

	
		std::cout << "Data Construction for algorithm is started....." << std::endl;
		// allocate host memory
		Node *h_graph_nodes = (Node*) malloc(sizeof(Node) * no_of_nodes);
		int *h_graph_mask = (int*) malloc(sizeof(int) * no_of_nodes);
		int *h_updating_graph_mask = (int*) malloc(sizeof(int) * no_of_nodes);
		int *h_graph_visited = (int*) malloc(sizeof(int) * no_of_nodes);
  
		// initalize the memory
		boost::graph_traits<Graph>::vertex_iterator vi, vi_end;
		size_t j = 0;
		size_t degree = 0;

		edge_list_size = 2 * num_edges(*g); //undirected

		//std::cout << "EDGE SIZE = " << edge_list_size << std::endl;
		int* h_graph_edges = (int*) malloc(sizeof(int) * edge_list_size);
		boost::graph_traits<Graph>::out_edge_iterator ei, ei_end;
		size_t i = 0;
		for (tie(vi, vi_end) = vertices(*g); vi != vi_end; ++vi) {
			h_graph_nodes[j].starting = degree;
			h_graph_nodes[j].no_of_edges = out_degree(*vi, *g);

			for (tie(ei, ei_end) = out_edges(*vi,*g); ei != ei_end; ++ei) {
				h_graph_edges[i] = target(*ei, *g);
				//std::cout << "The edge is : " << *ei << " target is :" << h_graph_edges[i] << std::endl;
				++i;
			}
		
			degree += h_graph_nodes[j].no_of_edges;
			//std::cout <<" starting for node "<< j <<"  "<< h_graph_nodes[j].starting 
				//<< "  number_of edges  " << h_graph_nodes[j].no_of_edges << std::endl;
			h_graph_mask[j]=false;
			h_updating_graph_mask[j]=false;
			h_graph_visited[j]=false;
			++j;
		}

		int source = *(vertices(*g).first);

		//set the source node as true in the mask
		h_graph_mask[source] = true;
		h_graph_visited[source] = true;
	/*
		edge_list_size = num_edges(*g);
		int* h_graph_edges = (int*) malloc(sizeof(int)*edge_list_size);

		boost::graph_traits<Graph>::edge_iterator ei, ei_end;
		size_t i = 0;
		for (tie(ei, ei_end) = out_edges(*vi,*g); ei != ei_end; ++ei) {
			h_graph_edges[i] = target(*ei, *g);
			std::cout << "The edge is : " << *ei << " target is :" << h_graph_edges[i] << std::endl;
			++i;
		}
		*/
		std::cout << "Data preparation for algorithm has been finished...." << std::endl;

		std::cout << "BFS algorithm is started on GPU device...." << std::endl;
	
		////Opencl variables.
		cl_device_id device_id = NULL;
		cl_context context = NULL;
		cl_command_queue command_queue = NULL;
		cl_mem d_graph_nodes = NULL;
		cl_mem d_graph_edges = NULL;
		cl_mem d_graph_mask = NULL;
		cl_mem d_updating_graph_mask = NULL;
		cl_mem d_graph_visited = NULL;
		cl_mem d_cost = NULL;
		cl_mem d_over = NULL;
		cl_program program = NULL;
		cl_kernel kernel = NULL;
		cl_program program2 = NULL;
		cl_kernel kernel2 = NULL;
		cl_platform_id platform_id = NULL;
		cl_uint ret_num_devices = 0;
		cl_uint ret_num_platforms = 0;
		cl_int ret = 0;

		FILE *fp;
		char fileName[] = "kernel.cl";
		char *source_str;
		size_t source_size;
	
		// Load the source code containing the kernel
		fp = fopen(fileName, "r");
		if (!fp) {
			fprintf(stderr, "Failed to load kernel.cl");
			exit(1);
			}
		source_str = (char*)malloc(MAX_SOURCE_SIZE);
		source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
		fclose(fp);
		

		FILE *fp2;
		char fileName2[] = "kernel2.cl";
		char *source_str2;
		size_t source_size2;

		// Load the source code containing the kernel
		fp2 = fopen(fileName2, "r");
		if (!fp2) {
			fprintf(stderr, "Failed to load kernel2.cl");
			exit(1);
			}
		source_str2 = (char*)malloc(MAX_SOURCE_SIZE);
		source_size2 = fread(source_str2, 1, MAX_SOURCE_SIZE, fp2);
		fclose(fp2);
	
		/* Get Platform and Device Info */
		ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
		if (CL_SUCCESS != ret) {
			fprintf(stderr, "Failed to get platform id !!!");
			exit(1);
		}

		/* Gets the all GPU devices */
		ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);
	
		/* Create OpenCL context */
		context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
	
		/* Create Command Queue */
		command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
	
		/* Create Memory Buffer for Nodes */
		d_graph_nodes = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY, sizeof(Node) * no_of_nodes,
								h_graph_nodes, &ret);
		//ret = clEnqueueWriteBuffer(command_queue, d_graph_nodes, CL_TRUE, 0, 
			//sizeof(Node) * no_of_nodes, h_graph_nodes, 0, NULL, NULL);
	/*
		ret = clEnqueueReadBuffer(command_queue, d_graph_nodes, CL_TRUE, 0, sizeof(Node) * no_of_nodes, h_graph_nodes, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"%d) graph node :%d\n",i, h_graph_nodes[i].starting);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
		*

		/* Create Memory Buffer for Edges */
		d_graph_edges = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY, sizeof(int) * edge_list_size, 
											h_graph_edges, &ret);
		//ret = clEnqueueWriteBuffer(command_queue, d_graph_edges, CL_TRUE, 0, 
			//sizeof(int) * edge_list_size, h_graph_edges, 0, NULL, NULL);
		/*	
		ret = clEnqueueReadBuffer(command_queue, d_graph_edges, CL_TRUE, 0, sizeof(int) * edge_list_size, 
			h_graph_edges, 0, NULL, NULL);
		
		//Store the result into a file
		FILE *fpo = fopen("result.txt","w");
		for(int i = 0; i < edge_list_size; i++) {
			fprintf(fpo,"%d) graph edge :%d\n", i, h_graph_edges[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
		*/
			
		/* Create Memory Buffer for masks */
		d_graph_mask = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, 
										sizeof(int) * no_of_nodes, h_graph_mask, &ret);
		//ret = clEnqueueWriteBuffer(command_queue, d_graph_mask, CL_TRUE, 0, 
			//sizeof(int) * no_of_nodes, h_graph_mask, 0, NULL, NULL);
	
	/*	ret = clEnqueueReadBuffer(command_queue, d_graph_mask, CL_TRUE, 0, 
					sizeof(int) * no_of_nodes, h_graph_mask, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"%d) graph mask:%d\n",i, h_graph_mask[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
		*/

		/* Create Memory Buffer for updating masks */
		d_updating_graph_mask = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, 
							sizeof(int) * no_of_nodes, h_updating_graph_mask, &ret);
		//ret = clEnqueueWriteBuffer(command_queue, d_updating_graph_mask, CL_TRUE, 0, 
			//sizeof(int) * no_of_nodes, h_updating_graph_mask, 0, NULL, NULL);
		/*
		ret = clEnqueueReadBuffer(command_queue, d_updating_graph_mask, CL_TRUE, 0, 
								sizeof(int) * no_of_nodes, h_updating_graph_mask, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"%d) graph updating mask:%d\n",i, h_updating_graph_mask[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
		*/
		/* Create Memory Buffer for visited */
		d_graph_visited = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, 
									sizeof(int) * no_of_nodes, h_graph_visited, &ret);

		//ret = clEnqueueWriteBuffer(command_queue, d_graph_visited, CL_TRUE, 0, 
			//sizeof(int) * no_of_nodes, h_graph_visited, 0, NULL, NULL);
	/*	
		ret = clEnqueueReadBuffer(command_queue, d_graph_visited, CL_TRUE, 0, 
			sizeof(int) * no_of_nodes, h_graph_visited, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"%d) graph visited:%d\n",i, h_graph_visited[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
		*/
	
		// allocate mem for the result on host side
		int* h_cost = (int*) malloc(sizeof(int) * no_of_nodes);
		for(int i = 0; i < no_of_nodes; i++) {
			h_cost[i] = -1;
		}
		h_cost[source] = 0;

		/* Create Memory Buffer for cost */
		d_cost = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sizeof(int) * no_of_nodes, h_cost, &ret);
		//ret = clEnqueueWriteBuffer(command_queue, d_cost, CL_TRUE, 0, 
			//sizeof(int) * no_of_nodes, h_cost, 0, NULL, NULL);
		/*
		ret = clEnqueueReadBuffer(command_queue, d_cost, CL_TRUE, 0, 
			sizeof(int) * no_of_nodes, h_cost, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"%d) graph cost:%d\n", i, h_cost[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");

		printf("Copied Everything to GPU memory\n");
		*/
	
		/* Create Kernel Program from the source */
		program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
			(const size_t *)&source_size, &ret);
		program2 = clCreateProgramWithSource(context, 1, (const char **)&source_str2,
			(const size_t *)&source_size2, &ret);
	
		/* Build Kernel Program */
		ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
		ret = clBuildProgram(program2, 1, &device_id, NULL, NULL, NULL);

		/* Create OpenCL Kernel */
		kernel = clCreateKernel(program, "Kernel", &ret);
		kernel2 = clCreateKernel(program2, "Kernel2", &ret);


	
		//make a bool to check if the execution is over
		d_over = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), 0, &ret);

		/* Set OpenCL Kernel Parameters */
		ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_graph_nodes);
		ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_graph_edges);
		ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&d_graph_mask);
		ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&d_updating_graph_mask);
		ret = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *)&d_graph_visited);
		ret = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *)&d_cost);
		ret = clSetKernelArg(kernel, 6, sizeof(int), (const void *)&no_of_nodes);

		ret = clSetKernelArg(kernel2, 0, sizeof(cl_mem), (void *)&d_graph_mask);
		ret = clSetKernelArg(kernel2, 1, sizeof(cl_mem), (void *)&d_updating_graph_mask);
		ret = clSetKernelArg(kernel2, 2, sizeof(cl_mem), (void *)&d_graph_visited);
		ret = clSetKernelArg(kernel2, 3, sizeof(cl_mem), (void *)&d_over);
		ret = clSetKernelArg(kernel2, 4, sizeof(int), (const void *)&no_of_nodes);
				// copy result from device to host

		/* Execute OpenCL Kernel */
	
		int k = 0;
		printf("Start traversing the tree\n");
		clock_t init, final;
		init = clock();

		//bool stop;
		int stop;
		//Call the Kernel untill all the elements of Frontier are not false
		do
		{
 		//for (int i = 0; i < 1; ++i) {
			//if no thread changes this value then the loop stops
 			//stop = false;
			stop = 0;
			ret = clEnqueueWriteBuffer(command_queue, d_over, CL_TRUE, 0, sizeof(int), &stop, 0, NULL, NULL);

			cl_event events[2];
			
			ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &no_of_nodes, NULL,  NULL, NULL, &events[0]);

			ret = clEnqueueNDRangeKernel(command_queue, kernel2, 1, NULL, &no_of_nodes, NULL, 1, &events[0], &events[1] );
		
			ret = clEnqueueReadBuffer(command_queue, d_over, CL_TRUE, 0, sizeof(int), &stop, 2, events, NULL);
			k++;
		}
		while( 1 == stop);

		final = clock() - init;
		std::cout << "The duration of execution is : "  << (double)final/((double)CLOCKS_PER_SEC)  << "  seconds"<< std::endl;
	
		printf("Kernel Executed %d times\n", k);	
		
		ret = clEnqueueReadBuffer(command_queue, d_cost, CL_TRUE, 0, 
			sizeof(int) * no_of_nodes, h_cost, 0, NULL, NULL);

		//Store the result into a file
		FILE *fpo = fopen("result.txt", "w");
		for(int i = 0; i < no_of_nodes; i++) {
			fprintf(fpo,"MinDistance(0, %d) = %d\n", i, h_cost[i]);
		}
		fclose(fpo);
		printf("Result stored in result.txt\n");
	
		std::cout << "BFS algorithm has been finished...." << std::endl;
	
		/* Finalization */
		ret = clFlush(command_queue);
		ret = clFinish(command_queue);
		ret = clReleaseKernel(kernel);
		ret = clReleaseProgram(program);
		ret = clReleaseMemObject(d_graph_nodes);
		ret = clReleaseMemObject(d_graph_edges);
		ret = clReleaseMemObject(d_graph_mask);
		ret = clReleaseMemObject(d_updating_graph_mask);
		ret = clReleaseMemObject(d_graph_visited);
		ret = clReleaseMemObject(d_cost);
		ret = clReleaseMemObject(d_over);
		ret = clReleaseCommandQueue(command_queue);
		ret = clReleaseContext(context);
		//free(source_str);	
	}
};

// This is the constructor of a class that has been exported.
// see Csharp_C.h for the class definition
CCsharp_C::CCsharp_C()
{
	return;
}
